<!-- HTML header for doxygen 1.8.3.1-->
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<meta name="generator" content="Doxygen 1.8.4"/>
<title>CUB: cub::BlockRadixSort&lt; Key, BLOCK_THREADS, ITEMS_PER_THREAD, Value, RADIX_BITS, MEMOIZE_OUTER_SCAN, INNER_SCAN_ALGORITHM, SMEM_CONFIG &gt; Class Template Reference</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
  $(document).ready(function() { searchBox.OnSelectItem(0); });
</script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
<link href="extra_stylesheet.css" rel="stylesheet" type="text/css"/>
<link rel="shortcut icon" href="favicon.ico" type="image/x-icon" />
<script type="text/javascript">
  var _gaq = _gaq || [];
  _gaq.push(['_setAccount', 'UA-38890655-1']);
  _gaq.push(['_trackPageview']);
  (function() {
    var ga = document.createElement('script'); ga.type = 'text/javascript'; ga.async = true;
    ga.src = ('https:' == document.location.protocol ? 'https://ssl' : 'http://www') + '.google-analytics.com/ga.js';
    var s = document.getElementsByTagName('script')[0]; s.parentNode.insertBefore(ga, s);
  })();
</script>
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
 <tbody>
 <tr style="height: 56px;">
  <td style="padding-left: 0.5em;">
   <div id="projectname">CUB
   </div>
  </td>
 </tr>
 </tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.4 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
  <div id="navrow1" class="tabs">
    <ul class="tablist">
      <li><a href="index.html"><span>Main&#160;Page</span></a></li>
      <li><a href="modules.html"><span>Modules</span></a></li>
      <li class="current"><a href="annotated.html"><span>Classes</span></a></li>
      <li>
        <div id="MSearchBox" class="MSearchBoxInactive">
        <span class="left">
          <img id="MSearchSelect" src="search/mag_sel.png"
               onmouseover="return searchBox.OnSearchSelectShow()"
               onmouseout="return searchBox.OnSearchSelectHide()"
               alt=""/>
          <input type="text" id="MSearchField" value="Search" accesskey="S"
               onfocus="searchBox.OnSearchFieldFocus(true)" 
               onblur="searchBox.OnSearchFieldFocus(false)" 
               onkeyup="searchBox.OnSearchFieldChange(event)"/>
          </span><span class="right">
            <a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
          </span>
        </div>
      </li>
    </ul>
  </div>
  <div id="navrow2" class="tabs2">
    <ul class="tablist">
      <li><a href="annotated.html"><span>Class&#160;List</span></a></li>
      <li><a href="classes.html"><span>Class&#160;Index</span></a></li>
    </ul>
  </div>
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
     onmouseover="return searchBox.OnSearchSelectShow()"
     onmouseout="return searchBox.OnSearchSelectHide()"
     onkeydown="return searchBox.OnSearchSelectKey(event)">
<a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(0)"><span class="SelectionMark">&#160;</span>All</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(1)"><span class="SelectionMark">&#160;</span>Classes</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(2)"><span class="SelectionMark">&#160;</span>Namespaces</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(3)"><span class="SelectionMark">&#160;</span>Files</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(4)"><span class="SelectionMark">&#160;</span>Functions</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(5)"><span class="SelectionMark">&#160;</span>Variables</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(6)"><span class="SelectionMark">&#160;</span>Typedefs</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(7)"><span class="SelectionMark">&#160;</span>Enumerations</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(8)"><span class="SelectionMark">&#160;</span>Enumerator</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(9)"><span class="SelectionMark">&#160;</span>Groups</a></div>

<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0" 
        name="MSearchResults" id="MSearchResults">
</iframe>
</div>

<div id="nav-path" class="navpath">
  <ul>
<li class="navelem"><a class="el" href="namespacecub.html">cub</a></li><li class="navelem"><a class="el" href="classcub_1_1_block_radix_sort.html">BlockRadixSort</a></li>  </ul>
</div>
</div><!-- top -->
<div class="header">
  <div class="summary">
<a href="#nested-classes">Classes</a> &#124;
<a href="classcub_1_1_block_radix_sort-members.html">List of all members</a>  </div>
  <div class="headertitle">
<div class="title">cub::BlockRadixSort&lt; Key, BLOCK_THREADS, ITEMS_PER_THREAD, Value, RADIX_BITS, MEMOIZE_OUTER_SCAN, INNER_SCAN_ALGORITHM, SMEM_CONFIG &gt; Class Template Reference<div class="ingroups"><a class="el" href="group___block_module.html">Block-wide</a></div></div>  </div>
</div><!--header-->
<div class="contents">
<a name="details" id="details"></a><h2 class="groupheader">Detailed description</h2>
<div class="textblock"><h3>template&lt;
    typename Key, 
    int BLOCK_THREADS, 
    int ITEMS_PER_THREAD, 
    typename Value = NullType, 
    int RADIX_BITS = 4, 
    bool MEMOIZE_OUTER_SCAN = (CUB_PTX_ARCH &gt;= 350) ? true : false, 
    BlockScanAlgorithm INNER_SCAN_ALGORITHM = BLOCK_SCAN_WARP_SCANS, 
    cudaSharedMemConfig SMEM_CONFIG = cudaSharedMemBankSizeFourByte&gt;<br/>
class cub::BlockRadixSort&lt; Key, BLOCK_THREADS, ITEMS_PER_THREAD, Value, RADIX_BITS, MEMOIZE_OUTER_SCAN, INNER_SCAN_ALGORITHM, SMEM_CONFIG &gt;</h3>

<p>The <a class="el" href="classcub_1_1_block_radix_sort.html" title="The cub::BlockRadixSort class provides collective methods for sorting items partitioned across a CUDA...">cub::BlockRadixSort</a> class provides <a href="index.html#sec0"><em>collective</em></a> methods for sorting items partitioned across a CUDA thread block using a radix sorting method. </p>
<div class="image">
<img src="sorting_logo.png" alt="sorting_logo.png"/>
<div class="caption">
.</div></div>
 <dl class="section user"><dt>Overview</dt><dd>The <a href="http://en.wikipedia.org/wiki/Radix_sort"><em>radix sorting method</em></a> arranges items into ascending order. It relies upon a positional representation for keys, i.e., each key is comprised of an ordered sequence of symbols (e.g., digits, characters, etc.) specified from least-significant to most-significant. For a given input sequence of keys and a set of rules specifying a total ordering of the symbolic alphabet, the radix sorting method produces a lexicographic ordering of those keys.</dd></dl>
<dl class="section user"><dt></dt><dd><a class="el" href="classcub_1_1_block_radix_sort.html" title="The cub::BlockRadixSort class provides collective methods for sorting items partitioned across a CUDA...">BlockRadixSort</a> can sort all of the built-in C++ numeric primitive types, e.g.: <code>unsigned char</code>, <code>int</code>, <code>double</code>, etc. Within each key, the implementation treats fixed-length bit-sequences of <code>RADIX_BITS</code> as radix digit places. Although the direct radix sorting method can only be applied to unsigned integral types, <a class="el" href="classcub_1_1_block_radix_sort.html" title="The cub::BlockRadixSort class provides collective methods for sorting items partitioned across a CUDA...">BlockRadixSort</a> is able to sort signed and floating-point types via simple bit-wise transformations that ensure lexicographic key ordering.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">Key</td><td>Key type </td></tr>
    <tr><td class="paramname">BLOCK_THREADS</td><td>The thread block size in threads </td></tr>
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td>The number of items per thread </td></tr>
    <tr><td class="paramname">Value</td><td><b>[optional]</b> Value type (default: <a class="el" href="structcub_1_1_null_type.html" title="A simple &quot;NULL&quot; marker type. ">cub::NullType</a>) </td></tr>
    <tr><td class="paramname">RADIX_BITS</td><td><b>[optional]</b> The number of radix bits per digit place (default: 4 bits) </td></tr>
    <tr><td class="paramname">MEMOIZE_OUTER_SCAN</td><td><b>[optional]</b> Whether or not to buffer outer raking scan partials to incur fewer shared memory reads at the expense of higher register pressure (default: true for architectures SM35 and newer, false otherwise). </td></tr>
    <tr><td class="paramname">INNER_SCAN_ALGORITHM</td><td><b>[optional]</b> The <a class="el" href="namespacecub.html#abec44bba36037c547e7e84906d0d23ab" title="BlockScanAlgorithm enumerates alternative algorithms for cub::BlockScan to compute a parallel prefix ...">cub::BlockScanAlgorithm</a> algorithm to use (default: <a class="el" href="namespacecub.html#abec44bba36037c547e7e84906d0d23aba7f51e58246eb53f1a97bd1bc8c0f400f">cub::BLOCK_SCAN_WARP_SCANS</a>) </td></tr>
    <tr><td class="paramname">SMEM_CONFIG</td><td><b>[optional]</b> Shared memory bank mode (default: <code>cudaSharedMemBankSizeFourByte</code>)</td></tr>
  </table>
  </dd>
</dl>
<dl class="section user"><dt>A Simple Example</dt><dd>Every thread in the block uses the <a class="el" href="classcub_1_1_block_radix_sort.html" title="The cub::BlockRadixSort class provides collective methods for sorting items partitioned across a CUDA...">BlockRadixSort</a> class by first specializing the <a class="el" href="classcub_1_1_block_radix_sort.html" title="The cub::BlockRadixSort class provides collective methods for sorting items partitioned across a CUDA...">BlockRadixSort</a> type, then instantiating an instance with parameters for communication, and finally invoking collective member functions. </dd></dl>
<dl class="section user"><dt></dt><dd>The code snippet below illustrates a sort of 512 integer keys that are partitioned in a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> across 128 threads where each thread owns 4 consecutive items. </dd></dl>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockRadixSort for 128 threads owning 4 integer items each</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_radix_sort.html" title="The cub::BlockRadixSort class provides collective methods for sorting items partitioned across a CUDA...">cub::BlockRadixSort&lt;int, 128, 4&gt;</a> <a class="code" href="classcub_1_1_block_radix_sort.html#ac7c896d6549960400b7710a262532e0e" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockRadixSort</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockRadixSort</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockRadixSort::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain a segment of consecutive items that are blocked across threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_keys[4];</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively sort the keys</span></div>
<div class="line">    <a class="code" href="classcub_1_1_block_radix_sort.html#ac7c896d6549960400b7710a262532e0e" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockRadixSort</a>(temp_storage).Sort(thread_keys);</div>
<div class="line"></div>
<div class="line">    ...</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_keys</code> across the block of threads is <code>{ [0,511,1,510], [2,509,3,508], [4,507,5,506], ..., [254,257,255,256] }</code>. The corresponding output <code>thread_keys</code> in those threads will be <code>{ [0,1,2,3], [4,5,6,7], [8,9,10,11], ..., [508,509,510,511] }</code>. </dd></dl>

<p>Definition at line <a class="el" href="block__radix__sort_8cuh_source.html#l00122">122</a> of file <a class="el" href="block__radix__sort_8cuh_source.html">block_radix_sort.cuh</a>.</p>
</div><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="nested-classes"></a>
Classes</h2></td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_block_radix_sort_1_1_temp_storage.html">TempStorage</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">The operations exposed by <a class="el" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">BlockScan</a> require a temporary memory allocation of this nested type for thread communication. This opaque storage can be allocated directly using the <code>__shared__</code> keyword. Alternatively, it can be aliased to externally allocated memory (shared or global) or <code>union</code>'d with other storage allocation types to facilitate memory reuse.  <a href="structcub_1_1_block_radix_sort_1_1_temp_storage.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="pub-methods"></a>
Public Methods</h2></td></tr>
<tr><td colspan="2"><div class="groupHeader">Collective constructors</div></td></tr>
<tr class="memitem:ac7c896d6549960400b7710a262532e0e"><td class="memItemLeft" align="right" valign="top"><a class="anchor" id="ac7c896d6549960400b7710a262532e0e"></a>
__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_radix_sort.html#ac7c896d6549960400b7710a262532e0e">BlockRadixSort</a> ()</td></tr>
<tr class="memdesc:ac7c896d6549960400b7710a262532e0e"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor for 1D thread blocks using a private static allocation of shared memory as temporary storage. Threads are identified using <code>threadIdx.x</code>. <br/></td></tr>
<tr class="separator:ac7c896d6549960400b7710a262532e0e"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:af6840b09ce6bb530fffa56ed75e4529c"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_radix_sort.html#af6840b09ce6bb530fffa56ed75e4529c">BlockRadixSort</a> (<a class="el" href="structcub_1_1_block_radix_sort_1_1_temp_storage.html">TempStorage</a> &amp;temp_storage)</td></tr>
<tr class="memdesc:af6840b09ce6bb530fffa56ed75e4529c"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage. Threads are identified using <code>threadIdx.x</code>.  <a href="#af6840b09ce6bb530fffa56ed75e4529c">More...</a><br/></td></tr>
<tr class="separator:af6840b09ce6bb530fffa56ed75e4529c"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a311b6e5932dfea96e1844abaf428bfc3"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_radix_sort.html#a311b6e5932dfea96e1844abaf428bfc3">BlockRadixSort</a> (int linear_tid)</td></tr>
<tr class="memdesc:a311b6e5932dfea96e1844abaf428bfc3"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor using a private static allocation of shared memory as temporary storage. Each thread is identified using the supplied linear thread identifier.  <a href="#a311b6e5932dfea96e1844abaf428bfc3">More...</a><br/></td></tr>
<tr class="separator:a311b6e5932dfea96e1844abaf428bfc3"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a34b8e018887026786b02e8ac6556bd3a"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_radix_sort.html#a34b8e018887026786b02e8ac6556bd3a">BlockRadixSort</a> (<a class="el" href="structcub_1_1_block_radix_sort_1_1_temp_storage.html">TempStorage</a> &amp;temp_storage, int linear_tid)</td></tr>
<tr class="memdesc:a34b8e018887026786b02e8ac6556bd3a"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor using the specified memory allocation as temporary storage. Each thread is identified using the supplied linear thread identifier.  <a href="#a34b8e018887026786b02e8ac6556bd3a">More...</a><br/></td></tr>
<tr class="separator:a34b8e018887026786b02e8ac6556bd3a"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Sorting (blocked arrangements)</div></td></tr>
<tr class="memitem:ad04651e04e555350bd250c4681a97644"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_radix_sort.html#ad04651e04e555350bd250c4681a97644">Sort</a> (Key(&amp;keys)[ITEMS_PER_THREAD], int begin_bit=0, int end_bit=sizeof(Key)*8)</td></tr>
<tr class="memdesc:ad04651e04e555350bd250c4681a97644"><td class="mdescLeft">&#160;</td><td class="mdescRight">Performs a block-wide radix sort over a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of keys.  <a href="#ad04651e04e555350bd250c4681a97644">More...</a><br/></td></tr>
<tr class="separator:ad04651e04e555350bd250c4681a97644"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:aac3d9424388b0cb6da360624e627b61e"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_radix_sort.html#aac3d9424388b0cb6da360624e627b61e">Sort</a> (Key(&amp;keys)[ITEMS_PER_THREAD], Value(&amp;values)[ITEMS_PER_THREAD], int begin_bit=0, int end_bit=sizeof(Key)*8)</td></tr>
<tr class="memdesc:aac3d9424388b0cb6da360624e627b61e"><td class="mdescLeft">&#160;</td><td class="mdescRight">Performs a block-wide radix sort across a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of keys and values.  <a href="#aac3d9424388b0cb6da360624e627b61e">More...</a><br/></td></tr>
<tr class="separator:aac3d9424388b0cb6da360624e627b61e"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Sorting (blocked arrangement -&gt; striped arrangement)</div></td></tr>
<tr class="memitem:a4be0c8d355a7fc75f8b405a776d884d3"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_radix_sort.html#a4be0c8d355a7fc75f8b405a776d884d3">SortBlockedToStriped</a> (Key(&amp;keys)[ITEMS_PER_THREAD], int begin_bit=0, int end_bit=sizeof(Key)*8)</td></tr>
<tr class="memdesc:a4be0c8d355a7fc75f8b405a776d884d3"><td class="mdescLeft">&#160;</td><td class="mdescRight">Performs a radix sort across a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of keys, leaving them in a <a href="index.html#sec5sec4"><em>striped arrangement</em></a>.  <a href="#a4be0c8d355a7fc75f8b405a776d884d3">More...</a><br/></td></tr>
<tr class="separator:a4be0c8d355a7fc75f8b405a776d884d3"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a9f707b1592f7354db0fba76a42afbc84"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_radix_sort.html#a9f707b1592f7354db0fba76a42afbc84">SortBlockedToStriped</a> (Key(&amp;keys)[ITEMS_PER_THREAD], Value(&amp;values)[ITEMS_PER_THREAD], int begin_bit=0, int end_bit=sizeof(Key)*8)</td></tr>
<tr class="memdesc:a9f707b1592f7354db0fba76a42afbc84"><td class="mdescLeft">&#160;</td><td class="mdescRight">Performs a radix sort across a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of keys and values, leaving them in a <a href="index.html#sec5sec4"><em>striped arrangement</em></a>.  <a href="#a9f707b1592f7354db0fba76a42afbc84">More...</a><br/></td></tr>
<tr class="separator:a9f707b1592f7354db0fba76a42afbc84"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table>
<h2 class="groupheader">Constructor &amp; Destructor Documentation</h2>
<a class="anchor" id="af6840b09ce6bb530fffa56ed75e4529c"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename Key , int BLOCK_THREADS, int ITEMS_PER_THREAD, typename Value  = NullType, int RADIX_BITS = 4, bool MEMOIZE_OUTER_SCAN = (CUB_PTX_ARCH &gt;= 350) ? true : false, BlockScanAlgorithm INNER_SCAN_ALGORITHM = BLOCK_SCAN_WARP_SCANS, cudaSharedMemConfig SMEM_CONFIG = cudaSharedMemBankSizeFourByte&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ <a class="el" href="classcub_1_1_block_radix_sort.html">cub::BlockRadixSort</a>&lt; Key, BLOCK_THREADS, ITEMS_PER_THREAD, Value, RADIX_BITS, MEMOIZE_OUTER_SCAN, INNER_SCAN_ALGORITHM, SMEM_CONFIG &gt;::<a class="el" href="classcub_1_1_block_radix_sort.html">BlockRadixSort</a> </td>
          <td>(</td>
          <td class="paramtype"><a class="el" href="structcub_1_1_block_radix_sort_1_1_temp_storage.html">TempStorage</a> &amp;&#160;</td>
          <td class="paramname"><em>temp_storage</em>)</td><td></td>
          <td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage. Threads are identified using <code>threadIdx.x</code>. </p>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">temp_storage</td><td>Reference to memory allocation having layout type <a class="el" href="structcub_1_1_block_radix_sort_1_1_temp_storage.html" title="The operations exposed by BlockScan require a temporary memory allocation of this nested type for thr...">TempStorage</a> </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__radix__sort_8cuh_source.html#l00201">201</a> of file <a class="el" href="block__radix__sort_8cuh_source.html">block_radix_sort.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a311b6e5932dfea96e1844abaf428bfc3"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename Key , int BLOCK_THREADS, int ITEMS_PER_THREAD, typename Value  = NullType, int RADIX_BITS = 4, bool MEMOIZE_OUTER_SCAN = (CUB_PTX_ARCH &gt;= 350) ? true : false, BlockScanAlgorithm INNER_SCAN_ALGORITHM = BLOCK_SCAN_WARP_SCANS, cudaSharedMemConfig SMEM_CONFIG = cudaSharedMemBankSizeFourByte&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ <a class="el" href="classcub_1_1_block_radix_sort.html">cub::BlockRadixSort</a>&lt; Key, BLOCK_THREADS, ITEMS_PER_THREAD, Value, RADIX_BITS, MEMOIZE_OUTER_SCAN, INNER_SCAN_ALGORITHM, SMEM_CONFIG &gt;::<a class="el" href="classcub_1_1_block_radix_sort.html">BlockRadixSort</a> </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>)</td><td></td>
          <td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Collective constructor using a private static allocation of shared memory as temporary storage. Each thread is identified using the supplied linear thread identifier. </p>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td>A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__radix__sort_8cuh_source.html#l00212">212</a> of file <a class="el" href="block__radix__sort_8cuh_source.html">block_radix_sort.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a34b8e018887026786b02e8ac6556bd3a"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename Key , int BLOCK_THREADS, int ITEMS_PER_THREAD, typename Value  = NullType, int RADIX_BITS = 4, bool MEMOIZE_OUTER_SCAN = (CUB_PTX_ARCH &gt;= 350) ? true : false, BlockScanAlgorithm INNER_SCAN_ALGORITHM = BLOCK_SCAN_WARP_SCANS, cudaSharedMemConfig SMEM_CONFIG = cudaSharedMemBankSizeFourByte&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ <a class="el" href="classcub_1_1_block_radix_sort.html">cub::BlockRadixSort</a>&lt; Key, BLOCK_THREADS, ITEMS_PER_THREAD, Value, RADIX_BITS, MEMOIZE_OUTER_SCAN, INNER_SCAN_ALGORITHM, SMEM_CONFIG &gt;::<a class="el" href="classcub_1_1_block_radix_sort.html">BlockRadixSort</a> </td>
          <td>(</td>
          <td class="paramtype"><a class="el" href="structcub_1_1_block_radix_sort_1_1_temp_storage.html">TempStorage</a> &amp;&#160;</td>
          <td class="paramname"><em>temp_storage</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Collective constructor using the specified memory allocation as temporary storage. Each thread is identified using the supplied linear thread identifier. </p>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">temp_storage</td><td>Reference to memory allocation having layout type <a class="el" href="structcub_1_1_block_radix_sort_1_1_temp_storage.html" title="The operations exposed by BlockScan require a temporary memory allocation of this nested type for thr...">TempStorage</a> </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td><b>[optional]</b> A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__radix__sort_8cuh_source.html#l00223">223</a> of file <a class="el" href="block__radix__sort_8cuh_source.html">block_radix_sort.cuh</a>.</p>

</div>
</div>
<h2 class="groupheader">Member Function Documentation</h2>
<a class="anchor" id="ad04651e04e555350bd250c4681a97644"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename Key , int BLOCK_THREADS, int ITEMS_PER_THREAD, typename Value  = NullType, int RADIX_BITS = 4, bool MEMOIZE_OUTER_SCAN = (CUB_PTX_ARCH &gt;= 350) ? true : false, BlockScanAlgorithm INNER_SCAN_ALGORITHM = BLOCK_SCAN_WARP_SCANS, cudaSharedMemConfig SMEM_CONFIG = cudaSharedMemBankSizeFourByte&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_radix_sort.html">cub::BlockRadixSort</a>&lt; Key, BLOCK_THREADS, ITEMS_PER_THREAD, Value, RADIX_BITS, MEMOIZE_OUTER_SCAN, INNER_SCAN_ALGORITHM, SMEM_CONFIG &gt;::Sort </td>
          <td>(</td>
          <td class="paramtype">Key(&amp;)&#160;</td>
          <td class="paramname"><em>keys</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>begin_bit</em> = <code>0</code>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>end_bit</em> = <code>sizeof(Key)&#160;*&#160;8</code>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Performs a block-wide radix sort over a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of keys. </p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates a sort of 512 integer keys that are partitioned in a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> across 128 threads where each thread owns 4 consecutive keys. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockRadixSort for 128 threads owning 4 integer keys each</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_radix_sort.html" title="The cub::BlockRadixSort class provides collective methods for sorting items partitioned across a CUDA...">cub::BlockRadixSort&lt;int, 128, 4&gt;</a> <a class="code" href="classcub_1_1_block_radix_sort.html#ac7c896d6549960400b7710a262532e0e" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockRadixSort</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockRadixSort</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockRadixSort::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain a segment of consecutive items that are blocked across threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_keys[4];</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively sort the keys</span></div>
<div class="line">    <a class="code" href="classcub_1_1_block_radix_sort.html#ac7c896d6549960400b7710a262532e0e" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockRadixSort</a>(temp_storage).Sort(thread_keys);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_keys</code> across the block of threads is <code>{ [0,511,1,510], [2,509,3,508], [4,507,5,506], ..., [254,257,255,256] }</code>. The corresponding output <code>thread_keys</code> in those threads will be <code>{ [0,1,2,3], [4,5,6,7], [8,9,10,11], ..., [508,509,510,511] }</code>. </dd></dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in,out]</td><td class="paramname">keys</td><td>Keys to sort </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">begin_bit</td><td><b>[optional]</b> The beginning (least-significant) bit index needed for key comparison </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">end_bit</td><td><b>[optional]</b> The past-the-end (most-significant) bit index needed for key comparison </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__radix__sort_8cuh_source.html#l00273">273</a> of file <a class="el" href="block__radix__sort_8cuh_source.html">block_radix_sort.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="aac3d9424388b0cb6da360624e627b61e"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename Key , int BLOCK_THREADS, int ITEMS_PER_THREAD, typename Value  = NullType, int RADIX_BITS = 4, bool MEMOIZE_OUTER_SCAN = (CUB_PTX_ARCH &gt;= 350) ? true : false, BlockScanAlgorithm INNER_SCAN_ALGORITHM = BLOCK_SCAN_WARP_SCANS, cudaSharedMemConfig SMEM_CONFIG = cudaSharedMemBankSizeFourByte&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_radix_sort.html">cub::BlockRadixSort</a>&lt; Key, BLOCK_THREADS, ITEMS_PER_THREAD, Value, RADIX_BITS, MEMOIZE_OUTER_SCAN, INNER_SCAN_ALGORITHM, SMEM_CONFIG &gt;::Sort </td>
          <td>(</td>
          <td class="paramtype">Key(&amp;)&#160;</td>
          <td class="paramname"><em>keys</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">Value(&amp;)&#160;</td>
          <td class="paramname"><em>values</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>begin_bit</em> = <code>0</code>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>end_bit</em> = <code>sizeof(Key)&#160;*&#160;8</code>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Performs a block-wide radix sort across a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of keys and values. </p>
<p><a class="el" href="classcub_1_1_block_radix_sort.html" title="The cub::BlockRadixSort class provides collective methods for sorting items partitioned across a CUDA...">BlockRadixSort</a> can only accommodate one associated tile of values. To "truck along" more than one tile of values, simply perform a key-value sort of the keys paired with a temporary value array that enumerates the key indices. The reordered indices can then be used as a gather-vector for exchanging other associated tile data through shared memory.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates a sort of 512 integer keys and values that are partitioned in a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> across 128 threads where each thread owns 4 consecutive pairs. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockRadixSort for 128 threads owning 4 integer keys and values each</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_radix_sort.html" title="The cub::BlockRadixSort class provides collective methods for sorting items partitioned across a CUDA...">cub::BlockRadixSort&lt;int, 128, 4, int&gt;</a> <a class="code" href="classcub_1_1_block_radix_sort.html#ac7c896d6549960400b7710a262532e0e" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockRadixSort</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockRadixSort</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockRadixSort::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain a segment of consecutive items that are blocked across threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_keys[4];</div>
<div class="line">    <span class="keywordtype">int</span> thread_values[4];</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively sort the keys and values among block threads</span></div>
<div class="line">    <a class="code" href="classcub_1_1_block_radix_sort.html#ac7c896d6549960400b7710a262532e0e" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockRadixSort</a>(temp_storage).Sort(thread_keys, thread_values);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_keys</code> across the block of threads is <code>{ [0,511,1,510], [2,509,3,508], [4,507,5,506], ..., [254,257,255,256] }</code>. The corresponding output <code>thread_keys</code> in those threads will be <code>{ [0,1,2,3], [4,5,6,7], [8,9,10,11], ..., [508,509,510,511] }</code>. </dd></dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in,out]</td><td class="paramname">keys</td><td>Keys to sort </td></tr>
    <tr><td class="paramdir">[in,out]</td><td class="paramname">values</td><td>Values to sort </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">begin_bit</td><td><b>[optional]</b> The beginning (least-significant) bit index needed for key comparison </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">end_bit</td><td><b>[optional]</b> The past-the-end (most-significant) bit index needed for key comparison </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__radix__sort_8cuh_source.html#l00358">358</a> of file <a class="el" href="block__radix__sort_8cuh_source.html">block_radix_sort.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a4be0c8d355a7fc75f8b405a776d884d3"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename Key , int BLOCK_THREADS, int ITEMS_PER_THREAD, typename Value  = NullType, int RADIX_BITS = 4, bool MEMOIZE_OUTER_SCAN = (CUB_PTX_ARCH &gt;= 350) ? true : false, BlockScanAlgorithm INNER_SCAN_ALGORITHM = BLOCK_SCAN_WARP_SCANS, cudaSharedMemConfig SMEM_CONFIG = cudaSharedMemBankSizeFourByte&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_radix_sort.html">cub::BlockRadixSort</a>&lt; Key, BLOCK_THREADS, ITEMS_PER_THREAD, Value, RADIX_BITS, MEMOIZE_OUTER_SCAN, INNER_SCAN_ALGORITHM, SMEM_CONFIG &gt;::SortBlockedToStriped </td>
          <td>(</td>
          <td class="paramtype">Key(&amp;)&#160;</td>
          <td class="paramname"><em>keys</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>begin_bit</em> = <code>0</code>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>end_bit</em> = <code>sizeof(Key)&#160;*&#160;8</code>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Performs a radix sort across a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of keys, leaving them in a <a href="index.html#sec5sec4"><em>striped arrangement</em></a>. </p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates a sort of 512 integer keys that are initially partitioned in a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> across 128 threads where each thread owns 4 consecutive keys. The final partitioning is striped. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockRadixSort for 128 threads owning 4 integer keys each</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_radix_sort.html" title="The cub::BlockRadixSort class provides collective methods for sorting items partitioned across a CUDA...">cub::BlockRadixSort&lt;int, 128, 4&gt;</a> <a class="code" href="classcub_1_1_block_radix_sort.html#ac7c896d6549960400b7710a262532e0e" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockRadixSort</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockRadixSort</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockRadixSort::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain a segment of consecutive items that are blocked across threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_keys[4];</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively sort the keys</span></div>
<div class="line">    <a class="code" href="classcub_1_1_block_radix_sort.html#ac7c896d6549960400b7710a262532e0e" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockRadixSort</a>(temp_storage).SortBlockedToStriped(thread_keys);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_keys</code> across the block of threads is <code>{ [0,511,1,510], [2,509,3,508], [4,507,5,506], ..., [254,257,255,256] }</code>. The corresponding output <code>thread_keys</code> in those threads will be <code>{ [0,128,256,384], [1,129,257,385], [2,130,258,386], ..., [127,255,383,511] }</code>. </dd></dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in,out]</td><td class="paramname">keys</td><td>Keys to sort </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">begin_bit</td><td><b>[optional]</b> The beginning (least-significant) bit index needed for key comparison </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">end_bit</td><td><b>[optional]</b> The past-the-end (most-significant) bit index needed for key comparison </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__radix__sort_8cuh_source.html#l00449">449</a> of file <a class="el" href="block__radix__sort_8cuh_source.html">block_radix_sort.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a9f707b1592f7354db0fba76a42afbc84"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename Key , int BLOCK_THREADS, int ITEMS_PER_THREAD, typename Value  = NullType, int RADIX_BITS = 4, bool MEMOIZE_OUTER_SCAN = (CUB_PTX_ARCH &gt;= 350) ? true : false, BlockScanAlgorithm INNER_SCAN_ALGORITHM = BLOCK_SCAN_WARP_SCANS, cudaSharedMemConfig SMEM_CONFIG = cudaSharedMemBankSizeFourByte&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_radix_sort.html">cub::BlockRadixSort</a>&lt; Key, BLOCK_THREADS, ITEMS_PER_THREAD, Value, RADIX_BITS, MEMOIZE_OUTER_SCAN, INNER_SCAN_ALGORITHM, SMEM_CONFIG &gt;::SortBlockedToStriped </td>
          <td>(</td>
          <td class="paramtype">Key(&amp;)&#160;</td>
          <td class="paramname"><em>keys</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">Value(&amp;)&#160;</td>
          <td class="paramname"><em>values</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>begin_bit</em> = <code>0</code>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>end_bit</em> = <code>sizeof(Key)&#160;*&#160;8</code>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Performs a radix sort across a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of keys and values, leaving them in a <a href="index.html#sec5sec4"><em>striped arrangement</em></a>. </p>
<p><a class="el" href="classcub_1_1_block_radix_sort.html" title="The cub::BlockRadixSort class provides collective methods for sorting items partitioned across a CUDA...">BlockRadixSort</a> can only accommodate one associated tile of values. To "truck along" more than one tile of values, simply perform a key-value sort of the keys paired with a temporary value array that enumerates the key indices. The reordered indices can then be used as a gather-vector for exchanging other associated tile data through shared memory.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates a sort of 512 integer keys and values that are initially partitioned in a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> across 128 threads where each thread owns 4 consecutive pairs. The final partitioning is striped. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockRadixSort for 128 threads owning 4 integer keys and values each</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_radix_sort.html" title="The cub::BlockRadixSort class provides collective methods for sorting items partitioned across a CUDA...">cub::BlockRadixSort&lt;int, 128, 4, int&gt;</a> <a class="code" href="classcub_1_1_block_radix_sort.html#ac7c896d6549960400b7710a262532e0e" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockRadixSort</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockRadixSort</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockRadixSort::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain a segment of consecutive items that are blocked across threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_keys[4];</div>
<div class="line">    <span class="keywordtype">int</span> thread_values[4];</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively sort the keys and values among block threads</span></div>
<div class="line">    <a class="code" href="classcub_1_1_block_radix_sort.html#ac7c896d6549960400b7710a262532e0e" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockRadixSort</a>(temp_storage).SortBlockedToStriped(thread_keys, thread_values);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_keys</code> across the block of threads is <code>{ [0,511,1,510], [2,509,3,508], [4,507,5,506], ..., [254,257,255,256] }</code>. The corresponding output <code>thread_keys</code> in those threads will be <code>{ [0,128,256,384], [1,129,257,385], [2,130,258,386], ..., [127,255,383,511] }</code>. </dd></dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in,out]</td><td class="paramname">keys</td><td>Keys to sort </td></tr>
    <tr><td class="paramdir">[in,out]</td><td class="paramname">values</td><td>Values to sort </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">begin_bit</td><td><b>[optional]</b> The beginning (least-significant) bit index needed for key comparison </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">end_bit</td><td><b>[optional]</b> The past-the-end (most-significant) bit index needed for key comparison </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__radix__sort_8cuh_source.html#l00541">541</a> of file <a class="el" href="block__radix__sort_8cuh_source.html">block_radix_sort.cuh</a>.</p>

</div>
</div>
<hr/>The documentation for this class was generated from the following file:<ul>
<li><a class="el" href="block__radix__sort_8cuh_source.html">block_radix_sort.cuh</a></li>
</ul>
</div><!-- contents -->
<!-- HTML footer for doxygen 1.8.3.1-->
<!-- start footer part -->
<hr class="footer"/><address class="footer"><small>
Generated on Fri Aug 23 2013 17:31:11 for CUB by &#160;<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/>
</a> 1.8.4
<br>
&copy; 2013 NVIDIA Corporation
</small></address>
</body>
</html>
